#include "../common.h"

//---------------------CUDA头文件----------------
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <cuda_device_runtime_api.h>
//---------------------CUDA头文件----------------


/**
 * CUDA kernel,在GPU上执行的函数。
 * 上千个线程都是执行这个函数，每个Thread根据全局id作为坐标来访问像素
 * @param src 类型是PtrStepSz<uchar3>，相当于是GpuMat的精简版
 */
__global__ void kernel(PtrStepSz<uchar3> src){
    int i = threadIdx.x + blockIdx.x * blockDim.x;  // 行，竖直方向的全局Thread id
    int j = threadIdx.y + blockIdx.y * blockDim.y;  // 列，水平方向的全局Thread id
//    printf("%d %d\n",i,j);                        //可用printf来debug
    if(i<src.rows && j<src.cols)                    //! 要判断是否越界！！！
        src(i,j) = make_uchar3(i%255,j%255,0);      // 访问一个像素的方式是src(行坐标,列坐标)
    __syncthreads();                                // 同步所有线程
}


int main() {
    const int rows = 16*50;
    const int cols = 16*60;
    if(getCudaEnabledDeviceCount()==0){
        cerr<<"此OpenCV编译的时候没有启用CUDA模块"<<endl;
        return -1;
    }
    // 定义一个纯黑GpuMat
    GpuMat gpuMat(rows,cols,CV_8UC3);
    // 每个block共有16*16个线程
    dim3 threadsPerBlock(16, 16);
    // 计算竖直需要多少个block
    uint block_num_vertical = (rows+threadsPerBlock.x-1)/threadsPerBlock.x;
    // 计算水平需要多少个block
    uint block_num_horizontal = (cols+threadsPerBlock.y-1)/threadsPerBlock.y;
    dim3 numBlocks(block_num_vertical, block_num_horizontal);
    // 调用kernel，gpuMat的类型会自动转换为PtrStepSz<uchar3>
    kernel<<<numBlocks,threadsPerBlock>>>(gpuMat);
    // 从显存把数据下载到内存
    Mat local;
    gpuMat.download(local);
    // 显示
    imshow("s",local);
    imwrite("../../s.jpg",local);
    waitKey(0);
    return 0;
}

